CUDA C++ Debugging: Safer GPU Kernel Programming (Generative AI Programming in C++) by Spuler David

Author:Spuler, David
Language: eng
Format: epub
Publisher: Aussie AI Labs
Published: 2024-10-15T00:00:00+00:00

Let’s examine these issues in turn.

Threads-per-Block Multiple of 32

The number of threads per block (aka the “block size”) should be a multiple of the warp size, which is 32 threads. Hence, it can be as low as 32, but commonly recommended block sizes in real-world kernels are often 256 or 512. The maximum permitted by CUDA is 1024 threads per block.

This is not good:

float v[54];

...

int blocks = 2;

mykernel<<<blocks, 27>>>(v,n); // BAD

This might actually work, but it’s very inefficient, and offends the sensibilities of any experienced CUDA programmer, for reasons discussed below.

But first, note that this is worse:

float v[54];

...

int blocks = 54;

mykernel<<<blocks, 1>>>(v,n); // BAD

If the threads per block is not 32, or a multiple of 32, there will be odd threads in a warp that aren’t properly utilized (or might be doing the wrong thing). The reason is that CUDA allows threads in “warps” that contain exactly 32 threads. With the threads per block of 27, there were 5 extra threads, and for 1 thread per block, there were 31 wasted threads. So, instead, you want something more like this:

float v[64];

int n = 64;

...

int blocks = 2;

mykernel<<<blocks,32>>>(v,n); // BETTER

Or you can do this:

float v[64];

int n = 64;

...

int blocks = 1;

mykernel<<<blocks,64>>>(v,n); // BETTER

CUDA can only schedule 32 threads (a warp) at a time, and if you try to schedule less than that, the rest of the threads in that warp still run, which is a bug or a slug, or are unavailable to run, which is a slug.

Too Few Blocks

You don’t always know the incoming size N of your vector data structure (well, actually, you often do in AI engines, because they have static dimensions, but anyway). Let’s try to generalize our computation of how many blocks with each having a fixed number of threads. There’s a few basic points:

Each block has the same number of threads (i.e., the threads-per-block)

You can’t run half a block (all its threads will run, even if you don’t need that many).

Download

Copyright Disclaimer:
This site does not store any files on its server. We only index and link to content provided by other sites. Please contact the content providers to delete copyright contents if any and email us, we'll remove relevant links or contents immediately.